You write custom CUDA kernels to replace pytorch operators in given architecture to get speedups. You have complete freedom to choose set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.

**SPECIAL INSTRUCTIONS FOR MANHATTAN + SWISH FUSION:**

When implementing Manhattan Distance + Swish fusion, you MUST implement the following optimized strategy:

1. **FUSION ARCHITECTURE**: Combine Swish activation and Manhattan distance computation in a single kernel:
   - Apply Swish activation to input tensor x first
   - Compute Manhattan distance between activated x and y
   - Eliminate intermediate tensor storage for maximum efficiency
   - Store both activated output and distance results

2. **FLOAT4 VECTORIZATION**: Use float4 vectorization for maximum memory bandwidth utilization:
   - Process 4 elements simultaneously using float4 loads/stores
   - Apply Swish activation to all 4 components in parallel
   - Compute Manhattan distance for all 4 components together
   - Handle remaining elements with scalar processing

3. **WARP-LEVEL OPTIMIZATION**: Use warp-level processing for maximum performance:
   - Each block processes one sample from the batch
   - Use 8 warps per block (256 threads) for optimal GPU utilization
   - Use __shfl_down_sync for efficient warp-level reduction of sums
   - Divide feature dimensions among warps for parallel processing

4. **MEMORY COALESCING**: Ensure efficient memory access patterns:
   - Use float4 vectorized loads for coalesced memory access
   - Store Swish results using float4 vectorized stores
   - Each thread processes multiple elements with stride pattern
   - Minimize global memory accesses through fusion

5. **EFFICIENT SUM REDUCTION**: Implement optimized sum reduction for Manhattan distance:
cpp
// Warp-level sum reduction
for (int offset = 16; offset > 0; offset /= 2) {
    warp_sum += __shfl_down_sync(0xffffffff, warp_sum, offset);
}

// Cross-warp reduction using shared memory
extern __shared__ float shared_data[];
if (lane_id == 0) {
    shared_data[warp_id] = warp_sum;
}
__syncthreads();



6. **SWISH FUSION**: Integrate Swish activation seamlessly with vectorization:
cpp
// Apply Swish activation to float4 vector
// Swish(x) = x * sigmoid(x) = x / (1 + exp(-x))
float4 x_activated;
x_activated.x = x_val.x / (1.0f + expf(-x_val.x));
x_activated.y = x_val.y / (1.0f + expf(-x_val.y));
x_activated.z = x_val.z / (1.0f + expf(-x_val.z));
x_activated.w = x_val.w / (1.0f + expf(-x_val.w));

// Store activated output
swish_vec[vec_idx] = x_activated;

// Compute Manhattan distance for all 4 components
warp_sum += fabsf(x_activated.x - y_val.x) + fabsf(x_activated.y - y_val.y) + 
           fabsf(x_activated.z - y_val.z) + fabsf(x_activated.w - y_val.w);



7. **SHARED MEMORY PATTERN**: Use efficient shared memory organization:
cpp
// For sum reduction across warps
extern __shared__ float shared_data[];
if (lane_id == 0) {
    shared_data[warp_id] = warp_sum;
}
__syncthreads();

// Final sum calculation
if (tid == 0) {
    float total_sum = 0.0f;
    int num_warps = blockDim.x / 32;
    for (int i = 0; i < num_warps; i++) {
        total_sum += shared_data[i];
    }
    distances[sample_idx] = total_sum;
}



8. **BLOCK CONFIGURATION**: Use optimal settings for vectorized processing:
   - Block size: 256 threads (8 warps)
   - Shared memory: 8 * sizeof(float) for warp reduction results
   - One block per sample for maximum parallelism
   - Elements per warp: (feature_dim + 8 - 1) / 8

9. **PRECISION REQUIREMENTS**: Ensure exact mathematical alignment:
   - Swish Activation: swish(x) = x * sigmoid(x) = x / (1 + exp(-x))
   - Manhattan Distance: Σ|swish(x) - y|
   - Use fabsf for absolute value computation
   - Use expf for exponential computation
   - Verify with torch.allclose(rtol=1e-03, atol=1e-6)

10. **FUNCTION SIGNATURE**: The main CUDA function must accept all parameters:
cpp
torch::Tensor manhattan_swish_cuda(
    torch::Tensor x,
    torch::Tensor y
)



11. **MATHEMATICAL FORMULAS**: Implement exact mathematical operations:
    - Swish Activation: swish(x) = x * sigmoid(x) = x / (1 + exp(-x))
    - Absolute Difference: abs_diff = |swish(x) - y|
    - Manhattan Distance: manhattan_dist = Σabs_diff

12. **PYTHON CALLING CONVENTION**: The ModelNew forward method must pass parameters correctly:
python
def forward(self, x, y):
    return self.manhattan_swish.manhattan_swish_cuda(x, y)



13. **OUTPUT REQUIREMENTS**: Generate both distances and activated outputs:
    - Primary output: Manhattan distances after Swish activation [batch_size]
    - Secondary output: Swish activated tensor [batch_size, feature_dim]
    - Both outputs must match PyTorch reference implementation exactly

14. **PERFORMANCE OPTIMIZATIONS**: Include advanced optimizations:
    - Use fast math optimizations (--use_fast_math)
    - Optimize for compute capability 8.0+ (sm_80)
    - Use -O3 optimization level
    - Avoid bank conflicts in shared memory access
    - Use efficient memory access patterns

15. **ALGORITHM CHOICE**: Prioritize the vectorized fused approach:
    - float4 vectorization is mandatory for this implementation
    - Do NOT implement scalar-only versions
    - The fusion must happen at the CUDA kernel level, not Python level
    - Eliminate all intermediate tensor storage

16. **BOUNDARY HANDLING**: Properly handle non-multiple-of-4 feature dimensions:
    - Use float4 for vectorized processing of main portion
    - Handle remaining elements with scalar processing
    - Ensure no memory access violations
    - Maintain mathematical correctness for all dimensions

17. **SWISH OPTIMIZATION**: Use the most efficient Swish implementation:
    - Swish(x) = x / (1 + exp(-x)) is more efficient than x * sigmoid(x)
    - Pre-compute denominator: 1.0f + expf(-x)
    - Avoid redundant sigmoid calculations
    - Use single division per element

Here's the target architecture to optimize:

python
import torch
import torch.nn as nn

class Model(nn.Module):
"""
Manhattan Distance implementation.
Computes the Manhattan distance (L1 distance) between two sets of vectors.
"""
def init(self):
super(Model, self).init()

def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    """
    Compute Manhattan distance between x and y.

    Args:
        x (torch.Tensor): First set of vectors [batch_size, feature_dim]
        y (torch.Tensor): Second set of vectors [batch_size, feature_dim]

    Returns:
        torch.Tensor: Manhattan distances [batch_size]
    """
    # Input validation
    if x.shape != y.shape:
        raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
    
    if x.dim() != 2:
        raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
    
    # Compute Manhattan distance: Σ|x_i - y_i|
    manhattan_dist = torch.sum(torch.abs(x - y), dim=1)
    
    return manhattan_dist

batch_size = 512
feature_dim = 512

def get_inputs():
# Generate two sets of vectors
x = torch.randn(batch_size, feature_dim)
y = torch.randn(batch_size, feature_dim)
return [x, y]

def get_init_inputs():
return [] # No special initialization inputs needed



**EXPECTED OUTPUT STRUCTURE**:
Generate two files:
1. `manhattan_swish_cudacode.py` - Contains ModelNew class with Manhattan+Swish fusion using pure CUDA
2. `manhattan_swish_torchcode.py` - Contains the reference PyTorch implementation with Swish fusion

**KEY REQUIREMENTS**:
- The CUDA implementation must use pure CUDA functions only
- Must implement Swish activation before Manhattan distance computation
- Must use float4 vectorization for maximum performance
- Must use warp-level optimization for maximum performance
- Must use efficient sum reduction algorithm
- Must handle arbitrary tensor shapes (not just fixed dimensions)
- Must maintain mathematical precision with PyTorch implementation
- Must use optimal block configuration (256 threads, 8 warps)
- Expected speedup: 1.6-2.3x over PyTorch baseline
- Must use fast math optimizations for better performance
- Must be robust and handle edge cases properly
- Must use only pure CUDA functions (no PyTorch internal functions)
- Must use fabsf for absolute value computation
- Must use expf for exponential computation
- Must implement exact mathematical formulas for Swish and Manhattan distance
- Must generate both distance and activated output tensors
- Must use shared memory efficiently for warp-level sum reduction
- Must ensure coalesced memory access patterns
- Must eliminate intermediate tensor storage for maximum fusion benefits
- Must implement the complete fusion in a single CUDA kernel
- Must use float4 vectorization as the primary optimization strategy
- Must use the most efficient Swish formula: swish(x) = x / (1 + exp(-x))
